//
//  openclconvolutioner.cpp
//  LowPassFilter_Standalone
//
//  Created by kadlubek47 on 16.04.2013.
//  Copyright (c) 2013 kadlubek47. All rights reserved.
//

#include "openclconvolutioner.h"
#include <sys/stat.h>
#include "GeneralExceptions.h"

OpenCLConvolutioner::OpenCLConvolutioner(): OpenCLProcessorBase("/convolute_kernel.cl") {
	createKernel();
}

void OpenCLConvolutioner::createKernel() {
    cl_int status;
    
    //-----------------------------------------------------
    // STEP 8: Create the kernel
    //-----------------------------------------------------
    
    // Use clCreateKernel() to create a kernel from the
    // vector addition function (named "vecadd")
    kernel_ = clCreateKernel(program_, "convolute", &status);
}

OpenCLConvolutioner::~OpenCLConvolutioner() {
    clReleaseMemObject(buffer_x_);
    clReleaseMemObject(buffer_y_);
    clReleaseMemObject(buffer_h_);
    clReleaseMemObject(buffer_M_);
}

void OpenCLConvolutioner::process(float_type* x, unsgn_integer_type x_size, float_type* h, unsgn_integer_type M, float_type* y) {
	setSizeOfDeviceBuffers(x_size, M);
	
    cl_int status;
	
	//-----------------------------------------------------
    // Write host data to device buffers
    //-----------------------------------------------------
    
    status = clEnqueueWriteBuffer(
                                  cmdQueue_,
                                  buffer_x_,
                                  CL_FALSE,
                                  0,
                                  x_size_ * sizeof(float_type),
                                  x,
                                  0,
                                  NULL,
                                  NULL);
    
    status = clEnqueueWriteBuffer(
                                  cmdQueue_,
                                  buffer_h_,
                                  CL_FALSE,
                                  0,
                                  h_size_ * sizeof(float_type),
                                  h,
                                  0,
                                  NULL,
                                  NULL);
    
    status = clEnqueueWriteBuffer(
                                  cmdQueue_,
                                  buffer_M_,
                                  CL_FALSE,
                                  0,
                                  sizeof(unsgn_integer_type),
                                  &M,
                                  0,
                                  NULL,
                                  NULL);

    //-----------------------------------------------------
    // Set the kernel arguments
    //-----------------------------------------------------

    status  = clSetKernelArg(
                             kernel_,
                             0,
                             sizeof(cl_mem),
                             &buffer_x_);
    
    status  = clSetKernelArg(
                             kernel_,
                             1,
                             512,
                             NULL);
	
	status |= clSetKernelArg(
                             kernel_,
                             2,
                             sizeof(cl_mem),
                             &buffer_h_ );
    status |= clSetKernelArg(
                             kernel_,
                             3,
                             sizeof(cl_mem),
                             &buffer_M_ );
    
    status |= clSetKernelArg(
                             kernel_,
                             4,
                             sizeof(cl_mem),
                             &buffer_y_ );
    
    
	//-----------------------------------------------------
    // Enqueue the kernel for execution
    //-----------------------------------------------------
    
	size_t globalWorkSize[1];
    // There are 'elements' work-items
	globalWorkSize[0] = y_size_;
    
    //@TODO
    //size_t workGroupSize[1];
    
    status = clEnqueueNDRangeKernel(
                                    cmdQueue_,
                                    kernel_,
                                    1,
                                    NULL,
                                    globalWorkSize,
                                    NULL,
                                    0,
                                    NULL,
                                    NULL);
    
    //-----------------------------------------------------
    // Read the output buffer back to the host
    //-----------------------------------------------------
    
    clEnqueueReadBuffer(
                        cmdQueue_,
                        buffer_y_,
                        CL_TRUE,
                        0,
                        y_size_*sizeof(float_type),
                        y,
                        0,
                        NULL,
                        NULL);
    
    clFlush(cmdQueue_);
}

void OpenCLConvolutioner::setSizeOfDeviceBuffers(unsgn_integer_type x_size, unsgn_integer_type M) {
    
    if( x_size==0 )
        throw WrongFunArgVal_Exception();
        
    if(x_size==x_size_ && M==M_)
        return;
       
    cl_int status;
    
	//-----------------------------------------------------
    // STEP 5: Create device buffers
    //-----------------------------------------------------
    
    if(x_size_ != 0 && buffersFirstTimeSet)
    {
		clReleaseMemObject(buffer_x_);
		clReleaseMemObject(buffer_y_);
	}
       
    x_size_=x_size;
    M_=M;
       
    buffer_x_ = clCreateBuffer(
                                context_,
                                CL_MEM_READ_ONLY,
                                x_size_ * sizeof(float_type),
                                NULL,
                                &status);
    if(!buffersFirstTimeSet) {
        
        h_size_ = 2*M + 1;
        buffer_h_ = clCreateBuffer(
                                   context_,
                                   CL_MEM_READ_ONLY,
                                   h_size_ * sizeof(float_type),
                                   NULL,
                                   &status);
        
        buffer_M_ = clCreateBuffer(
                                   context_,
                                   CL_MEM_READ_ONLY,
                                   1 * sizeof(unsgn_integer_type),
                                   NULL,
                                   &status);
        
    }
    
    y_size_ = x_size_ - 2*M_;
    buffer_y_ = clCreateBuffer(
                                context_,
                                CL_MEM_WRITE_ONLY,
                                y_size_ * sizeof(float_type),
                                NULL,
                                &status);
    
    buffersFirstTimeSet=true;
    
}


